使用CUDA进行编程，主要目的就是时间上加速。为此，如何计时必不可少。在CUDA中，我们可以使用CPU计时函数和GPU计时函数。对于CPU计时，我们在之前的文章(精确系统计时：秒、毫秒、微秒)中已经介绍在一般的C/C++编程中的计时方法。下面我们介绍在CUDA中如何计时：
CPU计时
CUDA中的核函数是异步执行的，即调用核函数后（而非等待其运行结束）就继续执行后面的语句。因此，使用CPU计时的时候，我们需要加上同步函数，这样才能得到核函数的运行时间，否则就是调用时间。下面给出一个简单的实例，由于代码简单，这里不再过多说明：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
| #include"cuda_runtime.h"
#include"device_launch_parameters.h"
#include<iostream>
#include<random>
#include<ctime>
using namespace std;
__global__ void kernel_function()
{
printf("Hello From GPU\n");
}
int main()
{
float esp_time_cpu;
clock_t start_cpu, stop_cpu;
kernel_function << <1, 10>> > ();
start_cpu = clock();
kernel_function<<<1,10>>> ();
cudaDeviceSynchronize();
stop_cpu = clock();
esp_time_cpu = (float)(stop_cpu - start_cpu) / CLOCKS_PER_SEC;
printf("The time by host:\t%f(ms)\n", esp_time_cpu);
cudaDeviceReset();
return 0;
}
|
需要注意的是：在计时前最好先warming up一下，即先把要计时的函数运行一遍。
GPU计时
这里我们可以使用CUDA提供的事件管理API来实现计时，具体可以参考NVIDIA官方的文档。具体实例如下：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
| #include"cuda_runtime.h"
#include"device_launch_parameters.h"
#include<iostream>
#include<random>
#include<ctime>
using namespace std;
__global__ void kernel_function()
{
printf("Hello From GPU\n");
}
Timing using GPU
int main()
{
cudaEvent_t start, stop;
float esp_time_gpu;
cudaEventCreate(&start);
cudaEventCreate(&stop);
kernel_function << <1, 10 >> > ();
cudaEventRecord(start, 0);
kernel_function << <1, 10 >> > ();
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&esp_time_gpu, start, stop);
printf("Time for the kernel: %f ms\n", esp_time_gpu);
cudaDeviceReset();
return 0;
}
|
CPU和GPU计时的结果分别如下：
注意：上述的结果不同，是因为有误差，我们可以通过多运行几次取平均值。
另外，上述CPU和GPU计时上的主要区别在于是否需要同步：CPU计时需要在核函数后调用同步函数。为此，当需要同步时，我们可以使用CPU计时；当不能同步时，我们使用GPU计时。当然，GPU计时也能用于需要同步的场景，此时的同步函数cudaDeviceSynchronize()需要放置在计时模块的后面。
